#include <stdio.h>
#include <stdlib.h>
#include <sys/types.h>
#include <sys/stat.h>
#include <fcntl.h>
#include <sys/mman.h>
#include <string.h>
#include <unistd.h>
#include <time.h>
#include <CL/cl.h>
#include <string>
#include <malloc.h>

#define BILLION 1000000000

using namespace std;
extern const char ** createProgram(string, int *);
extern const char * createProgramBinary(string, size_t *);

int main(void){

    struct timespec start, afBuild, end;
    int psc = 0;
    const char **ps = createProgram("select.cl",&psc);

    cl_int error = 0;
    cl_device_id device;

    cl_uint numP;
    clGetPlatformIDs(0,NULL,&numP);

    cl_platform_id * pid = new cl_platform_id[numP];
    clGetPlatformIDs(numP, pid, NULL);

    clGetDeviceIDs(pid[0],CL_DEVICE_TYPE_GPU,1,&device,NULL);

    cl_context context = clCreateContext(0,1,&device,NULL,NULL,&error);
    if(error != CL_SUCCESS){
        printf("Error creating context\n");
        exit(-1);
    }

    cl_command_queue_properties prop = 0;
    prop |= CL_QUEUE_PROFILING_ENABLE;

    cl_command_queue queue = clCreateCommandQueue(context, device, prop, &error);
    if(error != CL_SUCCESS){
        printf("Error creating command queue\n");
        exit(-1);
    }

    cl_program program = clCreateProgramWithSource(context, psc, (const char **)ps, 0, 0);
    if(error != CL_SUCCESS){
        printf("Error creating program\n");
        exit(-1);
    }

    clBuildProgram(program, 0, 0 , 0 , 0, 0);
    if(error != CL_SUCCESS){
            printf("Error building program\n");
            exit(-1);
    }

    int tupleNum = 1024*1024*64;
    int * input = (int *)memalign(256,sizeof(int)*tupleNum);
    int * output = (int *)memalign(256,sizeof(int)*tupleNum);

    for(int i=0;i<tupleNum;i++){
        input[i] = i;
    }

    memset(output,0,sizeof(int)*tupleNum);

    clock_gettime(CLOCK_REALTIME,&start);
    cl_mem column = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(int)*tupleNum, NULL, &error);
    //clEnqueueReadBuffer(queue, column, CL_TRUE, 0, sizeof(int)*(tupleNum), input,0,0,0);

    cl_event ndrEvt;
    cl_ulong startTime;
    cl_ulong endTime;

    cl_kernel kernel = clCreateKernel(program, "scan", 0);

    clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&column);
    clSetKernelArg(kernel, 1, sizeof(int), (void *)&tupleNum);

    size_t gl = 2048*256;
    size_t ll = 256;

    clEnqueueNDRangeKernel(queue, kernel, 1, 0, &gl,&ll,0,0,&ndrEvt);

    clWaitForEvents(1, &ndrEvt);
    clGetEventProfilingInfo(ndrEvt,
                                         CL_PROFILING_COMMAND_START,
                                         sizeof(cl_ulong),
                                         &startTime,
                                         0);

    clGetEventProfilingInfo(ndrEvt,
                                         CL_PROFILING_COMMAND_END,
                                         sizeof(cl_ulong),
                                         &endTime,
                                         0);

    printf("kernel %f\n",1e-6*(endTime - startTime));

    if(error != CL_SUCCESS){
                printf("prefix sum error %d\n",error);
                exit(-1);
     }

    
    clEnqueueReadBuffer(queue, column, CL_TRUE, 0, sizeof(int)*(tupleNum), output,0,0,&ndrEvt);

    clWaitForEvents(1, &ndrEvt);
    clGetEventProfilingInfo(ndrEvt,
                                         CL_PROFILING_COMMAND_START,
                                         sizeof(cl_ulong),
                                         &startTime,
                                         0);

    clGetEventProfilingInfo(ndrEvt,
                                         CL_PROFILING_COMMAND_END,
                                         sizeof(cl_ulong),
                                         &endTime,
                                         0);

    printf("DtoH transfer %f\n",1e-6*(endTime - startTime));

    error = clFinish(queue);
    clReleaseCommandQueue(queue);

    clock_gettime(CLOCK_REALTIME,&end);
    double timeE = (end.tv_sec -  start.tv_sec)* BILLION + end.tv_nsec - start.tv_nsec;
    printf("Wallclocktime: %lf\n", timeE/1000/1000);

}


